/**
 * __device__ __forceinline__ int GetThreadNum(void);
 * @brief Get Thread(sip) number
 * @param
 * @return Thread(sip) number
 */

/**
 * __device__ __forceinline__ int GetThreadIdx(void)
 * @brief Get global thread(sip) idx
 * @param
 * @return global thread(sip) idx
 */

/**
 * __device__ void dot_general_fp32(int lhs_addr, int rhs_addr, int M, int K, int N, int reduce_index, int reduce_cnt, int out_addr)
 * @brief
 * M == 16X
 * N = 32X
 * K = 32X
 * Semantic: [M,K] * [K,N] == [M,N]
 * DataFormat: [K/32,M,32] * [N/32,K,32] == [M,N]
 * @param lhs_addr lhs addr
 * @param rhs_addr rhs addr
 * @param M
 * @param K
 * @param M
 * @param reduce_index reduce index
 * @param reduce_cnt total reduce times
 * @param out_addr output addr
 */

// 实现设备端kernel函数代码
// 本算子完成 C(MxN) = A(MxK) . B(KxN) 的计算
// 线程在 M 维并行，每个线程在 K 和 N 维进行切分
__attribute__((global, cooperative)) void kernel_gemm(float *lhs, float *rhs, float *out, int M, int K, int N)
{
    printf("%d", M);
}

void GCU_GEMM(float *__restrict dev_lhs, float *__restrict dev_rhs, float *__restrict dev_out, const int m, const int k, const int n)
{
    static const size_t blocks = 1;
    static const size_t threads = 1;

    // 调用kernel
    kernel_gemm<<<blocks, threads>>>(dev_lhs, dev_rhs, dev_out, m, k, n);
}